Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Enable fp8e5m2fnuz type #3570

Merged
merged 17 commits into from
Nov 29, 2024
Merged

Enable fp8e5m2fnuz type #3570

merged 17 commits into from
Nov 29, 2024

Conversation

CharlieL7
Copy link
Collaborator

  • Enables the E5M2 FNUZ datatype so that we have full support for all the current FP8 types
  • E5M2 FNUZ will not be used directly in model but might be useful when converting from OCP -> FNUZ
  • Lots of updated files, but mostly expanding test cases

@CharlieL7 CharlieL7 added the FP8 issues related to FP8 implemenation label Oct 29, 2024
@CharlieL7 CharlieL7 self-assigned this Oct 29, 2024
@CharlieL7 CharlieL7 requested a review from causten as a code owner October 29, 2024 18:33
TedThemistokleous
TedThemistokleous reviewed Nov 4, 2024
View reviewed changes
test/verify/test_gemm_add_broadcast2.cpp
@@ -55,5 +55,6 @@ template struct test_gemm_add_broadcast2<migraphx::shape::float_type>;
// template struct test_gemm_add_broadcast2<migraphx::shape::half_type>; // fails with CK,
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we have a CK feature for this or is there a workaround?

TedThemistokleous
TedThemistokleous reviewed Nov 4, 2024
View reviewed changes
test/verify/test_gemm_add.cpp
@@ -58,3 +58,6 @@ struct test_gemm_add : verify_program<test_gemm_add<DType>>
template struct test_gemm_add<migraphx::shape::float_type>;
template struct test_gemm_add<migraphx::shape::half_type>;
// TODO template struct test_gemm_add<migraphx::shape::fp8e4m3fnuz_type>;
// TODO template struct test_gemm_add<migraphx::shape::fp8e5m2fnuz_type>;
// TODO template struct test_gemm_add<migraphx::shape::fp8e4m2fn_type>;
// TODO template struct test_gemm_add<migraphx::shape::fp8e5m2_type>;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why are these TODO?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

these don't work, umang made an issue about them earlier

TedThemistokleous
TedThemistokleous reviewed Nov 4, 2024
View reviewed changes
test/verify/test_gemm_transposea.cpp
@@ -46,6 +46,7 @@ struct test_gemm_transposea : verify_program<test_gemm_transposea<DType>>
template struct test_gemm_transposea<migraphx::shape::float_type>;
template struct test_gemm_transposea<migraphx::shape::half_type>;
template struct test_gemm_transposea<migraphx::shape::fp8e4m3fnuz_type>;
template struct test_gemm_transposea<migraphx::shape::fp8e5m2fnuz_type>;
// TODO need hipblaslt support
// template struct test_gemm_transposea<migraphx::shape::fp8e4m3fn_type>;
// template struct test_gemm_transposea<migraphx::shape::fp8e5m2_type>;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Anyway to workaround and run these tests without hipblastlt?

TedThemistokleous
TedThemistokleous reviewed Nov 4, 2024
View reviewed changes
Copy link
Collaborator

@TedThemistokleous TedThemistokleous left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Initial comments. Few questions, nothing concerning after you fixed the API compatability.

Just more following up on some of your comments with tests and things.

@pfultz2 pfultz2 mentioned this pull request Nov 9, 2024
Merged
@codecov Codecov
Copy link

codecov bot commented Nov 19, 2024 •
edited
Loading

Codecov Report

All modified and coverable lines are covered by tests ✅

Project coverage is 92.19%. Comparing base (da545d2) to head (e7f0d4b).
Report is 1 commits behind head on develop.

Additional details and impacted files 
@@           Coverage Diff            @@
##           develop    #3570   +/-   ##
========================================
  Coverage    92.19%   92.19%           
========================================
  Files          513      513           
  Lines        21633    21638    +5     
========================================
+ Hits         19945    19950    +5     
  Misses        1688     1688

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

@CharlieL7 CharlieL7 marked this pull request as draft November 22, 2024 18:46
@CharlieL7 CharlieL7 marked this pull request as ready for review November 25, 2024 20:19
@CharlieL7
Merge branch 'enable_fp8e5m2fnuz_type' of github.com:ROCmSoftwarePlat…
0195fff 
…form/AMDMIGraphX into enable_fp8e5m2fnuz_type
@migraphx-bot
Copy link
Collaborator

Test Batch Rate new
e7f0d4		 Rate old
da545d		 Diff Compare
torchvision-resnet50 64 3,257.68 3,255.84 0.06% ✅
torchvision-resnet50_fp16 64 6,998.19 6,991.64 0.09% ✅
torchvision-densenet121 32 2,436.13 2,431.85 0.18% ✅
torchvision-densenet121_fp16 32 4,093.27 4,070.27 0.56% ✅
torchvision-inceptionv3 32 1,627.63 1,628.09 -0.03% ✅
torchvision-inceptionv3_fp16 32 2,743.18 2,747.87 -0.17% ✅
cadene-inceptionv4 16 765.37 764.69 0.09% ✅
cadene-resnext64x4 16 810.93 806.89 0.50% ✅
slim-mobilenet 64 7,463.58 7,464.92 -0.02% ✅
slim-nasnetalarge 64 208.46 208.40 0.03% ✅
slim-resnet50v2 64 3,440.51 3,441.42 -0.03% ✅
bert-mrpc-onnx 8 1,145.13 1,145.08 0.00% ✅
bert-mrpc-tf 1 461.72 461.72 0.00% ✅
pytorch-examples-wlang-gru 1 422.34 429.74 -1.72% ✅
pytorch-examples-wlang-lstm 1 393.67 480.37 -18.05% 🔴
torchvision-resnet50_1 1 772.55 770.17 0.31% ✅
cadene-dpn92_1 1 416.14 403.49 3.14% 🔆
cadene-resnext101_1 1 382.82 381.97 0.22% ✅
onnx-taau-downsample 1 346.01 346.00 0.00% ✅
dlrm-criteoterabyte 1 33.32 33.31 0.02% ✅
dlrm-criteoterabyte_fp16 1 52.76 52.72 0.06% ✅
agentmodel 1 8,450.66 8,212.54 2.90% ✅
unet_fp16 2 58.82 58.71 0.18% ✅
resnet50v1_fp16 1 942.64 938.46 0.44% ✅
resnet50v1_int8 1 1,005.45 999.92 0.55% ✅
bert_base_cased_fp16 64 1,170.40 1,169.75 0.06% ✅
bert_large_uncased_fp16 32 363.16 363.16 -0.00% ✅
bert_large_fp16 1 200.44 200.11 0.16% ✅
distilgpt2_fp16 16 2,197.87 2,198.98 -0.05% ✅
yolov5s 1 534.90 533.15 0.33% ✅
tinyllama 1 43.62 43.41 0.49% ✅
vicuna-fastchat 1 174.11 173.49 0.36% ✅
whisper-tiny-encoder 1 418.39 417.39 0.24% ✅
whisper-tiny-decoder 1 424.68 427.99 -0.77% ✅

This build is not recommended to merge 🔴

@migraphx-bot
Copy link
Collaborator


     ✅ bert-mrpc-onnx: PASSED: MIGraphX meets tolerance

     ✅ bert-mrpc-tf: PASSED: MIGraphX meets tolerance

     ✅ pytorch-examples-wlang-gru: PASSED: MIGraphX meets tolerance

     ✅ pytorch-examples-wlang-lstm: PASSED: MIGraphX meets tolerance

     ✅ torchvision-resnet50_1: PASSED: MIGraphX meets tolerance

     ✅ cadene-dpn92_1: PASSED: MIGraphX meets tolerance

     ✅ cadene-resnext101_1: PASSED: MIGraphX meets tolerance

     ✅ dlrm-criteoterabyte: PASSED: MIGraphX meets tolerance

     ✅ agentmodel: PASSED: MIGraphX meets tolerance

     ✅ unet: PASSED: MIGraphX meets tolerance

     ✅ resnet50v1: PASSED: MIGraphX meets tolerance

     ✅ bert_base_cased_fp16: PASSED: MIGraphX meets tolerance

🔴bert_large_uncased_fp16: FAILED: MIGraphX is not within tolerance - check verbose output


     ✅ bert_large: PASSED: MIGraphX meets tolerance

     ✅ yolov5s: PASSED: MIGraphX meets tolerance

     ✅ tinyllama: PASSED: MIGraphX meets tolerance

     ✅ vicuna-fastchat: PASSED: MIGraphX meets tolerance

     ✅ whisper-tiny-encoder: PASSED: MIGraphX meets tolerance

     ✅ whisper-tiny-decoder: PASSED: MIGraphX meets tolerance

     ✅ distilgpt2_fp16: PASSED: MIGraphX meets tolerance

@TedThemistokleous TedThemistokleous added the high priority A PR with high priority for review and merging. label Nov 29, 2024
@TedThemistokleous
Copy link
Collaborator

Should get this in sooner than later @causten . Onnxrt has support for this as well form the looks of it so it would be nice to get his all in for fp8 support at once

@causten causten merged commit 35fd39f into develop Nov 29, 2024
43 of 45 checks passed
@causten causten deleted the enable_fp8e5m2fnuz_type branch November 29, 2024 15:54
shivadbhavsar pushed a commit that referenced this pull request Dec 18, 2024
@CharlieL7 @shivadbhavsar
Enable fp8e5m2fnuz type (#3570)
5e18b64
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@pfultz2 pfultz2 pfultz2 approved these changes

@TedThemistokleous TedThemistokleous TedThemistokleous approved these changes

@causten causten Awaiting requested review from causten causten is a code owner

Assignees

@CharlieL7 CharlieL7

Labels
FP8 issues related to FP8 implemenation high priority A PR with high priority for review and merging.
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
5 participants
@CharlieL7 @migraphx-bot @TedThemistokleous @pfultz2 @causten